Yes, it still appears to be broken: http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26504/steps/run%20tests/logs/stdio
On Thu, Jul 28, 2016 at 1:30 PM, Samuel F Antao via cfe-commits <cfe-commits@lists.llvm.org> wrote: > Hi Mike, > > I've already pushed r276981 and r276988 to fix those failures shortly after > I pushed the patch. After that I didn't receive any bot message complaining > about that code. > > Do you still see the issues? > > Thanks, > Samuel > > > ----- Original message ----- > From: Mike Aizatsky <aizat...@google.com> > To: Samuel F Antao/Watson/IBM@IBMUS, cfe-commits@lists.llvm.org > Cc: > Subject: Re: r276977 - [OpenMP] Codegen for use_device_ptr clause. > Date: Thu, Jul 28, 2016 4:20 PM > > Samuel, > > I think your change breaks the build: > > http://lab.llvm.org:8011/builders/sanitizer-windows/builds/26397 > > FAILED: > tools/clang/lib/CodeGen/CMakeFiles/clangCodeGen.dir/CGOpenMPRuntime.cpp.obj > C:\PROGRA~2\MICROS~1.0\VC\bin\AMD64_~1\cl.exe /nologo /TP > -DCLANG_ENABLE_ARCMT -DCLANG_ENABLE_OBJC_REWRITER > -DCLANG_ENABLE_STATIC_ANALYZER -DGTEST_HAS_RTTI=0 -DUNICODE > -D_CRT_NONSTDC_NO_DEPRECATE -D_CRT_NONSTDC_NO_WARNINGS > -D_CRT_SECURE_NO_DEPRECATE -D_CRT_SECURE_NO_WARNINGS > -D_DEBUG_POINTER_IMPL="" -D_GNU_SOURCE -D_HAS_EXCEPTIONS=0 > -D_SCL_SECURE_NO_DEPRECATE -D_SCL_SECURE_NO_WARNINGS -D_UNICODE > -D__STDC_CONSTANT_MACROS -D__STDC_FORMAT_MACROS -D__STDC_LIMIT_MACROS > -Itools\clang\lib\CodeGen > -IC:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen > -IC:\b\slave\sanitizer-windows\llvm\tools\clang\include > -Itools\clang\include -Iinclude -IC:\b\slave\sanitizer-windows\llvm\include > /DWIN32 /D_WINDOWS /W4 -wd4141 -wd4146 -wd4180 -wd4244 -wd4258 -wd4267 > -wd4291 -wd4345 -wd4351 -wd4355 -wd4456 -wd4457 -wd4458 -wd4459 -wd4503 > -wd4624 -wd4722 -wd4800 -wd4100 -wd4127 -wd4512 -wd4505 -wd4610 -wd4510 > -wd4702 -wd4245 -wd4706 -wd4310 -wd4701 -wd4703 -wd4389 -wd4611 -wd4805 > -wd4204 -wd4577 -wd4091 -wd4592 -wd4319 -wd4324 -w14062 -we4238 /Zc:inline > /Oi /Zc:rvalueCast /MD /O2 /Ob2 -UNDEBUG /EHs-c- /GR- /showIncludes > /Fotools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\CGOpenMPRuntime.cpp.obj > /Fdtools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\ /FS -c > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5482) > : error C3486: a parameter for a lambda cannot have a default argument > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is > not a type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) > : error C2065: 'Directive' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) > : error C2228: left of '.getClausesOfKind' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) > : error C2059: syntax error : ')' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5488) > : error C2143: syntax error : missing ';' before '<end Parse>' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is > not a type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) > : error C2065: 'Directive' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) > : error C2228: left of '.getClausesOfKind' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) > : error C2059: syntax error : ')' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5491) > : error C2143: syntax error : missing ';' before '<end Parse>' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5493) > : error C2143: syntax error : missing ')' before ';' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is > not a type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) > : error C2065: 'Directive' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) > : error C2228: left of '.getClausesOfKind' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) > : error C2059: syntax error : ')' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5494) > : error C2143: syntax error : missing ';' before '<end Parse>' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5496) > : error C2143: syntax error : missing ')' before ';' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::Directive' : is > not a type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) > : error C2065: 'Directive' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) > : error C2228: left of '.getClausesOfKind' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) > : error C2059: syntax error : ')' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5502) > : error C2143: syntax error : missing ';' before '<end Parse>' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503) > : error C2143: syntax error : missing ';' before '{' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5503) > : error C2143: syntax error : missing ')' before ';' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) > : error C2065: 'L' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) > : error C2228: left of '.second' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5504) > : error C2228: left of '.empty' must have class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) > : error C2065: 'L' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) > : error C2228: left of '.second' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) > : error C2228: left of '.back' must have class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5505) > : error C2228: left of '.getAssociatedDeclaration' must have > class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) > : error C2065: 'L' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) > : error C2228: left of '.second' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) > : error C2228: left of '.back' must have class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5507) > : error C2228: left of '.getAssociatedExpression' must have > class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5511) > : error C3536: 'IE': cannot be used before it is initialized > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5523) > : error C3536: 'IE': cannot be used before it is initialized > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5528) > : error C2044: illegal continue > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::CGF' : is not a > type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) > : error C2065: 'CGF' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) > : error C2228: left of '.EmitLoadOfLValue' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) > : error C2228: left of '.EmitLValue' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5535) > : error C3536: 'IE': cannot be used before it is initialized > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5536) > : error C2228: left of '.getScalarVal' must have class/struct/union > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) > : error C2327: '`anonymous-namespace'::MappableExprsHandler::CGF' : is not a > type name, static, or enumerator > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) > : error C2065: 'CGF' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5539) > : error C2228: left of '.SizeTy' must have class/struct/union > type is 'unknown-type' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : error C2143: syntax error : missing ')' before ';' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : error C2143: syntax error : missing ';' before ')' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : error C2065: '$S239' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : error C2065: '$S240' : undeclared identifier > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : error C2143: syntax error : missing ';' before '{' > C:\b\slave\sanitizer-windows\llvm\tools\clang\lib\CodeGen\CGOpenMPRuntime.cpp(5543) > : fatal error C1903: unable to recover from previous error(s); stopping > compilation > 18446743528493.273 [136/2/116] Building CXX object > tools\clang\lib\CodeGen\CMakeFiles\clangCodeGen.dir\CGOpenCLRuntime.cpp.obj > > On Thu, Jul 28, 2016 at 7:31 AM Samuel Antao via cfe-commits > <cfe-commits@lists.llvm.org> wrote: > > Author: sfantao > Date: Thu Jul 28 09:23:26 2016 > New Revision: 276977 > > URL: http://llvm.org/viewvc/llvm-project?rev=276977&view=rev > Log: > [OpenMP] Codegen for use_device_ptr clause. > > Summary: This patch adds support for the use_device_ptr clause. It includes > changes in SEMA that could not be tested without codegen, namely, the use of > the first private logic and mappable expressions support. > > Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev > > Subscribers: caomhin, cfe-commits > > Differential Revision: https://reviews.llvm.org/D22691 > > Added: > cfe/trunk/test/OpenMP/target_data_use_device_ptr_codegen.cpp > Modified: > cfe/trunk/include/clang/AST/OpenMPClause.h > cfe/trunk/lib/AST/OpenMPClause.cpp > cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp > cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h > cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp > cfe/trunk/lib/CodeGen/CodeGenFunction.h > cfe/trunk/lib/Sema/SemaOpenMP.cpp > cfe/trunk/lib/Serialization/ASTReaderStmt.cpp > cfe/trunk/lib/Serialization/ASTWriterStmt.cpp > > Modified: cfe/trunk/include/clang/AST/OpenMPClause.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/include/clang/AST/OpenMPClause.h (original) > +++ cfe/trunk/include/clang/AST/OpenMPClause.h Thu Jul 28 09:23:26 2016 > @@ -4228,50 +4228,153 @@ public: > /// 'use_device_ptr' with the variables 'a' and 'b'. > /// > class OMPUseDevicePtrClause final > - : public OMPVarListClause<OMPUseDevicePtrClause>, > - private llvm::TrailingObjects<OMPUseDevicePtrClause, Expr *> { > + : public OMPMappableExprListClause<OMPUseDevicePtrClause>, > + private llvm::TrailingObjects< > + OMPUseDevicePtrClause, Expr *, ValueDecl *, unsigned, > + OMPClauseMappableExprCommon::MappableComponent> { > friend TrailingObjects; > friend OMPVarListClause; > + friend OMPMappableExprListClause; > friend class OMPClauseReader; > - /// Build clause with number of variables \a N. > + > + /// Define the sizes of each trailing object array except the last one. > This > + /// is required for TrailingObjects to work properly. > + size_t numTrailingObjects(OverloadToken<Expr *>) const { > + return 3 * varlist_size(); > + } > + size_t numTrailingObjects(OverloadToken<ValueDecl *>) const { > + return getUniqueDeclarationsNum(); > + } > + size_t numTrailingObjects(OverloadToken<unsigned>) const { > + return getUniqueDeclarationsNum() + getTotalComponentListNum(); > + } > + > + /// Build clause with number of variables \a NumVars. > /// > /// \param StartLoc Starting location of the clause. > - /// \param LParenLoc Location of '('. > /// \param EndLoc Ending location of the clause. > - /// \param N Number of the variables in the clause. > - /// > - OMPUseDevicePtrClause(SourceLocation StartLoc, SourceLocation LParenLoc, > - SourceLocation EndLoc, unsigned N) > - : OMPVarListClause<OMPUseDevicePtrClause>(OMPC_use_device_ptr, > StartLoc, > - LParenLoc, EndLoc, N) {} > + /// \param NumVars Number of expressions listed in this clause. > + /// \param NumUniqueDeclarations Number of unique base declarations in > this > + /// clause. > + /// \param NumComponentLists Number of component lists in this clause. > + /// \param NumComponents Total number of expression components in the > clause. > + /// > + explicit OMPUseDevicePtrClause(SourceLocation StartLoc, > + SourceLocation LParenLoc, > + SourceLocation EndLoc, unsigned NumVars, > + unsigned NumUniqueDeclarations, > + unsigned NumComponentLists, > + unsigned NumComponents) > + : OMPMappableExprListClause(OMPC_use_device_ptr, StartLoc, LParenLoc, > + EndLoc, NumVars, NumUniqueDeclarations, > + NumComponentLists, NumComponents) {} > > - /// \brief Build an empty clause. > - /// > - /// \param N Number of variables. > + /// Build an empty clause. > /// > - explicit OMPUseDevicePtrClause(unsigned N) > - : OMPVarListClause<OMPUseDevicePtrClause>( > - OMPC_use_device_ptr, SourceLocation(), SourceLocation(), > - SourceLocation(), N) {} > + /// \param NumVars Number of expressions listed in this clause. > + /// \param NumUniqueDeclarations Number of unique base declarations in > this > + /// clause. > + /// \param NumComponentLists Number of component lists in this clause. > + /// \param NumComponents Total number of expression components in the > clause. > + /// > + explicit OMPUseDevicePtrClause(unsigned NumVars, > + unsigned NumUniqueDeclarations, > + unsigned NumComponentLists, > + unsigned NumComponents) > + : OMPMappableExprListClause(OMPC_use_device_ptr, SourceLocation(), > + SourceLocation(), SourceLocation(), > NumVars, > + NumUniqueDeclarations, NumComponentLists, > + NumComponents) {} > + > + /// Sets the list of references to private copies with initializers for > new > + /// private variables. > + /// \param VL List of references. > + void setPrivateCopies(ArrayRef<Expr *> VL); > + > + /// Gets the list of references to private copies with initializers for > new > + /// private variables. > + MutableArrayRef<Expr *> getPrivateCopies() { > + return MutableArrayRef<Expr *>(varlist_end(), varlist_size()); > + } > + ArrayRef<const Expr *> getPrivateCopies() const { > + return llvm::makeArrayRef(varlist_end(), varlist_size()); > + } > + > + /// Sets the list of references to initializer variables for new private > + /// variables. > + /// \param VL List of references. > + void setInits(ArrayRef<Expr *> VL); > + > + /// Gets the list of references to initializer variables for new private > + /// variables. > + MutableArrayRef<Expr *> getInits() { > + return MutableArrayRef<Expr *>(getPrivateCopies().end(), > varlist_size()); > + } > + ArrayRef<const Expr *> getInits() const { > + return llvm::makeArrayRef(getPrivateCopies().end(), varlist_size()); > + } > > public: > - /// Creates clause with a list of variables \a VL. > + /// Creates clause with a list of variables \a Vars. > /// > /// \param C AST context. > /// \param StartLoc Starting location of the clause. > - /// \param LParenLoc Location of '('. > /// \param EndLoc Ending location of the clause. > - /// \param VL List of references to the variables. > + /// \param Vars The original expression used in the clause. > + /// \param PrivateVars Expressions referring to private copies. > + /// \param Inits Expressions referring to private copy initializers. > + /// \param Declarations Declarations used in the clause. > + /// \param ComponentLists Component lists used in the clause. > /// > static OMPUseDevicePtrClause * > Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation > LParenLoc, > - SourceLocation EndLoc, ArrayRef<Expr *> VL); > - /// Creates an empty clause with the place for \a N variables. > + SourceLocation EndLoc, ArrayRef<Expr *> Vars, > + ArrayRef<Expr *> PrivateVars, ArrayRef<Expr *> Inits, > + ArrayRef<ValueDecl *> Declarations, > + MappableExprComponentListsRef ComponentLists); > + > + /// Creates an empty clause with the place for \a NumVars variables. > /// > /// \param C AST context. > - /// \param N The number of variables. > - /// > - static OMPUseDevicePtrClause *CreateEmpty(const ASTContext &C, unsigned > N); > + /// \param NumVars Number of expressions listed in the clause. > + /// \param NumUniqueDeclarations Number of unique base declarations in > this > + /// clause. > + /// \param NumComponentLists Number of unique base declarations in this > + /// clause. > + /// \param NumComponents Total number of expression components in the > clause. > + /// > + static OMPUseDevicePtrClause *CreateEmpty(const ASTContext &C, > + unsigned NumVars, > + unsigned NumUniqueDeclarations, > + unsigned NumComponentLists, > + unsigned NumComponents); > + > + typedef MutableArrayRef<Expr *>::iterator private_copies_iterator; > + typedef ArrayRef<const Expr *>::iterator private_copies_const_iterator; > + typedef llvm::iterator_range<private_copies_iterator> > private_copies_range; > + typedef llvm::iterator_range<private_copies_const_iterator> > + private_copies_const_range; > + > + private_copies_range private_copies() { > + return private_copies_range(getPrivateCopies().begin(), > + getPrivateCopies().end()); > + } > + private_copies_const_range private_copies() const { > + return private_copies_const_range(getPrivateCopies().begin(), > + getPrivateCopies().end()); > + } > + > + typedef MutableArrayRef<Expr *>::iterator inits_iterator; > + typedef ArrayRef<const Expr *>::iterator inits_const_iterator; > + typedef llvm::iterator_range<inits_iterator> inits_range; > + typedef llvm::iterator_range<inits_const_iterator> inits_const_range; > + > + inits_range inits() { > + return inits_range(getInits().begin(), getInits().end()); > + } > + inits_const_range inits() const { > + return inits_const_range(getInits().begin(), getInits().end()); > + } > > child_range children() { > return child_range(reinterpret_cast<Stmt **>(varlist_begin()), > > Modified: cfe/trunk/lib/AST/OpenMPClause.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/AST/OpenMPClause.cpp (original) > +++ cfe/trunk/lib/AST/OpenMPClause.cpp Thu Jul 28 09:23:26 2016 > @@ -732,22 +732,66 @@ OMPFromClause *OMPFromClause::CreateEmpt > NumComponentLists, NumComponents); > } > > -OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create(const ASTContext &C, > - SourceLocation > StartLoc, > - SourceLocation > LParenLoc, > - SourceLocation EndLoc, > - ArrayRef<Expr *> VL) { > - void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size())); > - OMPUseDevicePtrClause *Clause = > - new (Mem) OMPUseDevicePtrClause(StartLoc, LParenLoc, EndLoc, > VL.size()); > - Clause->setVarRefs(VL); > +void OMPUseDevicePtrClause::setPrivateCopies(ArrayRef<Expr *> VL) { > + assert(VL.size() == varlist_size() && > + "Number of private copies is not the same as the preallocated > buffer"); > + std::copy(VL.begin(), VL.end(), varlist_end()); > +} > + > +void OMPUseDevicePtrClause::setInits(ArrayRef<Expr *> VL) { > + assert(VL.size() == varlist_size() && > + "Number of inits is not the same as the preallocated buffer"); > + std::copy(VL.begin(), VL.end(), getPrivateCopies().end()); > +} > + > +OMPUseDevicePtrClause *OMPUseDevicePtrClause::Create( > + const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, > + SourceLocation EndLoc, ArrayRef<Expr *> Vars, ArrayRef<Expr *> > PrivateVars, > + ArrayRef<Expr *> Inits, ArrayRef<ValueDecl *> Declarations, > + MappableExprComponentListsRef ComponentLists) { > + unsigned NumVars = Vars.size(); > + unsigned NumUniqueDeclarations = > + getUniqueDeclarationsTotalNumber(Declarations); > + unsigned NumComponentLists = ComponentLists.size(); > + unsigned NumComponents = getComponentsTotalNumber(ComponentLists); > + > + // We need to allocate: > + // 3 x NumVars x Expr* - we have an original list expression for each > clause > + // list entry and an equal number of private copies and inits. > + // NumUniqueDeclarations x ValueDecl* - unique base declarations > associated > + // with each component list. > + // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify > the > + // number of lists for each unique declaration and the size of each > component > + // list. > + // NumComponents x MappableComponent - the total of all the components in > all > + // the lists. > + void *Mem = C.Allocate( > + totalSizeToAlloc<Expr *, ValueDecl *, unsigned, > + OMPClauseMappableExprCommon::MappableComponent>( > + 3 * NumVars, NumUniqueDeclarations, > + NumUniqueDeclarations + NumComponentLists, NumComponents)); > + > + OMPUseDevicePtrClause *Clause = new (Mem) OMPUseDevicePtrClause( > + StartLoc, LParenLoc, EndLoc, NumVars, NumUniqueDeclarations, > + NumComponentLists, NumComponents); > + > + Clause->setVarRefs(Vars); > + Clause->setPrivateCopies(PrivateVars); > + Clause->setInits(Inits); > + Clause->setClauseInfo(Declarations, ComponentLists); > return Clause; > } > > -OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty(const ASTContext > &C, > - unsigned N) { > - void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N)); > - return new (Mem) OMPUseDevicePtrClause(N); > +OMPUseDevicePtrClause *OMPUseDevicePtrClause::CreateEmpty( > + const ASTContext &C, unsigned NumVars, unsigned NumUniqueDeclarations, > + unsigned NumComponentLists, unsigned NumComponents) { > + void *Mem = C.Allocate( > + totalSizeToAlloc<Expr *, ValueDecl *, unsigned, > + OMPClauseMappableExprCommon::MappableComponent>( > + 3 * NumVars, NumUniqueDeclarations, > + NumUniqueDeclarations + NumComponentLists, NumComponents)); > + return new (Mem) OMPUseDevicePtrClause(NumVars, NumUniqueDeclarations, > + NumComponentLists, NumComponents); > } > > OMPIsDevicePtrClause *OMPIsDevicePtrClause::Create(const ASTContext &C, > > Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original) > +++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Jul 28 09:23:26 2016 > @@ -4981,6 +4981,9 @@ public: > /// map/privatization results in multiple arguments passed to the > runtime > /// library. > OMP_MAP_FIRST_REF = 0x20, > + /// \brief Signal that the runtime library has to return the device > pointer > + /// in the current position for the data being mapped. > + OMP_MAP_RETURN_PTR = 0x40, > /// \brief This flag signals that the reference being passed is a > pointer to > /// private data. > OMP_MAP_PRIVATE_PTR = 0x80, > @@ -4988,6 +4991,24 @@ public: > OMP_MAP_PRIVATE_VAL = 0x100, > }; > > + /// Class that associates information with a base pointer to be passed to > the > + /// runtime library. > + class BasePointerInfo { > + /// The base pointer. > + llvm::Value *Ptr = nullptr; > + /// The base declaration that refers to this device pointer, or null if > + /// there is none. > + const ValueDecl *DevPtrDecl = nullptr; > + > + public: > + BasePointerInfo(llvm::Value *Ptr, const ValueDecl *DevPtrDecl = > nullptr) > + : Ptr(Ptr), DevPtrDecl(DevPtrDecl) {} > + llvm::Value *operator*() const { return Ptr; } > + const ValueDecl *getDevicePtrDecl() const { return DevPtrDecl; } > + void setDevicePtrDecl(const ValueDecl *D) { DevPtrDecl = D; } > + }; > + > + typedef SmallVector<BasePointerInfo, 16> MapBaseValuesArrayTy; > typedef SmallVector<llvm::Value *, 16> MapValuesArrayTy; > typedef SmallVector<unsigned, 16> MapFlagsArrayTy; > > @@ -5129,7 +5150,7 @@ private: > void generateInfoForComponentList( > OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier, > OMPClauseMappableExprCommon::MappableExprComponentListRef Components, > - MapValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, > + MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, > MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, > bool IsFirstComponentList) const { > > @@ -5400,8 +5421,10 @@ public: > } > > /// \brief Generate all the base pointers, section pointers, sizes and > map > - /// types for the extracted mappable expressions. > - void generateAllInfo(MapValuesArrayTy &BasePointers, > + /// types for the extracted mappable expressions. Also, for each item > that > + /// relates with a device pointer, a pair of the relevant declaration and > + /// index where it occurs is appended to the device pointers info array. > + void generateAllInfo(MapBaseValuesArrayTy &BasePointers, > MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, > MapFlagsArrayTy &Types) const { > BasePointers.clear(); > @@ -5410,9 +5433,28 @@ public: > Types.clear(); > > struct MapInfo { > + /// Kind that defines how a device pointer has to be returned. > + enum ReturnPointerKind { > + // Don't have to return any pointer. > + RPK_None, > + // Pointer is the base of the declaration. > + RPK_Base, > + // Pointer is a member of the base declaration - 'this' > + RPK_Member, > + // Pointer is a reference and a member of the base declaration - > 'this' > + RPK_MemberReference, > + }; > OMPClauseMappableExprCommon::MappableExprComponentListRef Components; > - OpenMPMapClauseKind MapType; > - OpenMPMapClauseKind MapTypeModifier; > + OpenMPMapClauseKind MapType = OMPC_MAP_unknown; > + OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; > + ReturnPointerKind ReturnDevicePointer = RPK_None; > + MapInfo( > + OMPClauseMappableExprCommon::MappableExprComponentListRef > Components, > + OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier, > + ReturnPointerKind ReturnDevicePointer) > + : Components(Components), MapType(MapType), > + MapTypeModifier(MapTypeModifier), > + ReturnDevicePointer(ReturnDevicePointer) {} > }; > > // We have to process the component lists that relate with the same > @@ -5422,14 +5464,15 @@ public: > > // Helper function to fill the information map for the different > supported > // clauses. > - auto &&InfoGen = > - [&Info](const ValueDecl *D, > - OMPClauseMappableExprCommon::MappableExprComponentListRef > L, > - OpenMPMapClauseKind MapType, OpenMPMapClauseKind > MapModifier) { > - const ValueDecl *VD = > - D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr; > - Info[VD].push_back({L, MapType, MapModifier}); > - }; > + auto &&InfoGen = [&Info]( > + const ValueDecl *D, > + OMPClauseMappableExprCommon::MappableExprComponentListRef L, > + OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier, > + MapInfo::ReturnPointerKind ReturnDevicePointer = MapInfo::RPK_None) > { > + const ValueDecl *VD = > + D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr; > + Info[VD].push_back({L, MapType, MapModifier, ReturnDevicePointer}); > + }; > > for (auto *C : Directive.getClausesOfKind<OMPMapClause>()) > for (auto L : C->component_lists()) > @@ -5441,6 +5484,51 @@ public: > for (auto L : C->component_lists()) > InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown); > > + // Look at the use_device_ptr clause information and mark the existing > map > + // entries as such. If there is no map information for an entry in the > + // use_device_ptr list, we create one with map type 'alloc' and zero > size > + // section. It is the user fault if that was not mapped before. > + for (auto *C : Directive.getClausesOfKind<OMPUseDevicePtrClause>()) > + for (auto L : C->component_lists()) { > + assert(!L.second.empty() && "Not expecting empty list of > components!"); > + const ValueDecl *VD = L.second.back().getAssociatedDeclaration(); > + VD = cast<ValueDecl>(VD->getCanonicalDecl()); > + auto *IE = L.second.back().getAssociatedExpression(); > + // If the first component is a member expression, we have to look > into > + // 'this', which maps to null in the map of map information. > Otherwise > + // look directly for the information. > + auto It = Info.find(isa<MemberExpr>(IE) ? nullptr : VD); > + > + // We potentially have map information for this declaration > already. > + // Look for the first set of components that refer to it. > + if (It != Info.end()) { > + auto CI = std::find_if( > + It->second.begin(), It->second.end(), [VD](const MapInfo &MI) > { > + return MI.Components.back().getAssociatedDeclaration() == > VD; > + }); > + // If we found a map entry, signal that the pointer has to be > returned > + // and move on to the next declaration. > + if (CI != It->second.end()) { > + CI->ReturnDevicePointer = isa<MemberExpr>(IE) > + ? > (VD->getType()->isReferenceType() > + ? > MapInfo::RPK_MemberReference > + : MapInfo::RPK_Member) > + : MapInfo::RPK_Base; > + continue; > + } > + } > + > + // We didn't find any match in our map information - generate a > zero > + // size array section. > + llvm::Value *Ptr = > + CGF.EmitLoadOfLValue(CGF.EmitLValue(IE), SourceLocation()) > + .getScalarVal(); > + BasePointers.push_back({Ptr, VD}); > + Pointers.push_back(Ptr); > + Sizes.push_back(llvm::Constant::getNullValue(CGF.SizeTy)); > + Types.push_back(OMP_MAP_RETURN_PTR | OMP_MAP_FIRST_REF); > + } > + > for (auto &M : Info) { > // We need to know when we generate information for the first > component > // associated with a capture, because the mapping flags depend on it. > @@ -5448,9 +5536,35 @@ public: > for (MapInfo &L : M.second) { > assert(!L.Components.empty() && > "Not expecting declaration with no component lists."); > + > + // Remember the current base pointer index. > + unsigned CurrentBasePointersIdx = BasePointers.size(); > generateInfoForComponentList(L.MapType, L.MapTypeModifier, > L.Components, > BasePointers, Pointers, Sizes, Types, > IsFirstComponentList); > + > + // If this entry relates with a device pointer, set the relevant > + // declaration and add the 'return pointer' flag. > + if (IsFirstComponentList && > + L.ReturnDevicePointer != MapInfo::RPK_None) { > + // If the pointer is not the base of the map, we need to skip the > + // base. If it is a reference in a member field, we also need to > skip > + // the map of the reference. > + if (L.ReturnDevicePointer != MapInfo::RPK_Base) { > + ++CurrentBasePointersIdx; > + if (L.ReturnDevicePointer == MapInfo::RPK_MemberReference) > + ++CurrentBasePointersIdx; > + } > + assert(BasePointers.size() > CurrentBasePointersIdx && > + "Unexpected number of mapped base pointers."); > + > + auto *RelevantVD = > L.Components.back().getAssociatedDeclaration(); > + assert(RelevantVD && > + "No relevant declaration related with device pointer??"); > + > + > BasePointers[CurrentBasePointersIdx].setDevicePtrDecl(RelevantVD); > + Types[CurrentBasePointersIdx] |= OMP_MAP_RETURN_PTR; > + } > IsFirstComponentList = false; > } > } > @@ -5459,7 +5573,7 @@ public: > /// \brief Generate the base pointers, section pointers, sizes and map > types > /// associated to a given capture. > void generateInfoForCapture(const CapturedStmt::Capture *Cap, > - MapValuesArrayTy &BasePointers, > + MapBaseValuesArrayTy &BasePointers, > MapValuesArrayTy &Pointers, > MapValuesArrayTy &Sizes, > MapFlagsArrayTy &Types) const { > @@ -5496,12 +5610,12 @@ public: > > /// \brief Generate the default map information for a given capture \a > CI, > /// record field declaration \a RI and captured value \a CV. > - void generateDefaultMapInfo( > - const CapturedStmt::Capture &CI, const FieldDecl &RI, llvm::Value > *CV, > - MappableExprsHandler::MapValuesArrayTy &CurBasePointers, > - MappableExprsHandler::MapValuesArrayTy &CurPointers, > - MappableExprsHandler::MapValuesArrayTy &CurSizes, > - MappableExprsHandler::MapFlagsArrayTy &CurMapTypes) { > + void generateDefaultMapInfo(const CapturedStmt::Capture &CI, > + const FieldDecl &RI, llvm::Value *CV, > + MapBaseValuesArrayTy &CurBasePointers, > + MapValuesArrayTy &CurPointers, > + MapValuesArrayTy &CurSizes, > + MapFlagsArrayTy &CurMapTypes) { > > // Do the default mapping. > if (CI.capturesThis()) { > @@ -5510,15 +5624,14 @@ public: > const PointerType *PtrTy = > cast<PointerType>(RI.getType().getTypePtr()); > CurSizes.push_back(CGF.getTypeSize(PtrTy->getPointeeType())); > // Default map type. > - CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_TO | > - MappableExprsHandler::OMP_MAP_FROM); > + CurMapTypes.push_back(OMP_MAP_TO | OMP_MAP_FROM); > } else if (CI.capturesVariableByCopy()) { > CurBasePointers.push_back(CV); > CurPointers.push_back(CV); > if (!RI.getType()->isAnyPointerType()) { > // We have to signal to the runtime captures passed by value that > are > // not pointers. > - CurMapTypes.push_back(MappableExprsHandler::OMP_MAP_PRIVATE_VAL); > + CurMapTypes.push_back(OMP_MAP_PRIVATE_VAL); > CurSizes.push_back(CGF.getTypeSize(RI.getType())); > } else { > // Pointers are implicitly mapped with a zero size and no flags > @@ -5539,9 +5652,8 @@ public: > // default the value doesn't have to be retrieved. For an aggregate > // type, the default is 'tofrom'. > CurMapTypes.push_back(ElementType->isAggregateType() > - ? (MappableExprsHandler::OMP_MAP_TO | > - MappableExprsHandler::OMP_MAP_FROM) > - : MappableExprsHandler::OMP_MAP_TO); > + ? (OMP_MAP_TO | OMP_MAP_FROM) > + : OMP_MAP_TO); > > // If we have a capture by reference we may need to add the private > // pointer flag if the base declaration shows in some first-private > @@ -5551,7 +5663,7 @@ public: > } > // Every default map produces a single argument, so, it is always the > // first one. > - CurMapTypes.back() |= MappableExprsHandler::OMP_MAP_FIRST_REF; > + CurMapTypes.back() |= OMP_MAP_FIRST_REF; > } > }; > > @@ -5566,19 +5678,20 @@ enum OpenMPOffloadingReservedDeviceIDs { > /// offloading runtime library. If there is no map or capture information, > /// return nullptr by reference. > static void > -emitOffloadingArrays(CodeGenFunction &CGF, llvm::Value *&BasePointersArray, > - llvm::Value *&PointersArray, llvm::Value *&SizesArray, > - llvm::Value *&MapTypesArray, > - MappableExprsHandler::MapValuesArrayTy &BasePointers, > +emitOffloadingArrays(CodeGenFunction &CGF, > + MappableExprsHandler::MapBaseValuesArrayTy > &BasePointers, > MappableExprsHandler::MapValuesArrayTy &Pointers, > MappableExprsHandler::MapValuesArrayTy &Sizes, > - MappableExprsHandler::MapFlagsArrayTy &MapTypes) { > + MappableExprsHandler::MapFlagsArrayTy &MapTypes, > + CGOpenMPRuntime::TargetDataInfo &Info) { > auto &CGM = CGF.CGM; > auto &Ctx = CGF.getContext(); > > - BasePointersArray = PointersArray = SizesArray = MapTypesArray = nullptr; > + // Reset the array information. > + Info.clearArrayInfo(); > + Info.NumberOfPtrs = BasePointers.size(); > > - if (unsigned PointerNumVal = BasePointers.size()) { > + if (Info.NumberOfPtrs) { > // Detect if we have any capture size requiring runtime evaluation of > the > // size so that a constant array could be eventually used. > bool hasRuntimeEvaluationCaptureSize = false; > @@ -5588,14 +5701,14 @@ emitOffloadingArrays(CodeGenFunction &CG > break; > } > > - llvm::APInt PointerNumAP(32, PointerNumVal, /*isSigned=*/true); > + llvm::APInt PointerNumAP(32, Info.NumberOfPtrs, /*isSigned=*/true); > QualType PointerArrayType = > Ctx.getConstantArrayType(Ctx.VoidPtrTy, PointerNumAP, > ArrayType::Normal, > /*IndexTypeQuals=*/0); > > - BasePointersArray = > + Info.BasePointersArray = > CGF.CreateMemTemp(PointerArrayType, > ".offload_baseptrs").getPointer(); > - PointersArray = > + Info.PointersArray = > CGF.CreateMemTemp(PointerArrayType, ".offload_ptrs").getPointer(); > > // If we don't have any VLA types or other types that require runtime > @@ -5605,7 +5718,7 @@ emitOffloadingArrays(CodeGenFunction &CG > QualType SizeArrayType = Ctx.getConstantArrayType( > Ctx.getSizeType(), PointerNumAP, ArrayType::Normal, > /*IndexTypeQuals=*/0); > - SizesArray = > + Info.SizesArray = > CGF.CreateMemTemp(SizeArrayType, ".offload_sizes").getPointer(); > } else { > // We expect all the sizes to be constant, so we collect them to > create > @@ -5621,7 +5734,7 @@ emitOffloadingArrays(CodeGenFunction &CG > /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage, > SizesArrayInit, ".offload_sizes"); > > SizesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global); > - SizesArray = SizesArrayGbl; > + Info.SizesArray = SizesArrayGbl; > } > > // The map types are always constant so we don't need to generate code > to > @@ -5633,10 +5746,10 @@ emitOffloadingArrays(CodeGenFunction &CG > /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage, > MapTypesArrayInit, ".offload_maptypes"); > > MapTypesArrayGbl->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::Global); > - MapTypesArray = MapTypesArrayGbl; > + Info.MapTypesArray = MapTypesArrayGbl; > > - for (unsigned i = 0; i < PointerNumVal; ++i) { > - llvm::Value *BPVal = BasePointers[i]; > + for (unsigned i = 0; i < Info.NumberOfPtrs; ++i) { > + llvm::Value *BPVal = *BasePointers[i]; > if (BPVal->getType()->isPointerTy()) > BPVal = CGF.Builder.CreateBitCast(BPVal, CGM.VoidPtrTy); > else { > @@ -5645,11 +5758,15 @@ emitOffloadingArrays(CodeGenFunction &CG > BPVal = CGF.Builder.CreateIntToPtr(BPVal, CGM.VoidPtrTy); > } > llvm::Value *BP = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), > BasePointersArray, > - 0, i); > + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), > + Info.BasePointersArray, 0, i); > Address BPAddr(BP, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy)); > CGF.Builder.CreateStore(BPVal, BPAddr); > > + if (Info.requiresDevicePointerInfo()) > + if (auto *DevVD = BasePointers[i].getDevicePtrDecl()) > + Info.CaptureDeviceAddrMap.insert(std::make_pair(DevVD, BPAddr)); > + > llvm::Value *PVal = Pointers[i]; > if (PVal->getType()->isPointerTy()) > PVal = CGF.Builder.CreateBitCast(PVal, CGM.VoidPtrTy); > @@ -5659,14 +5776,15 @@ emitOffloadingArrays(CodeGenFunction &CG > PVal = CGF.Builder.CreateIntToPtr(PVal, CGM.VoidPtrTy); > } > llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.VoidPtrTy, PointerNumVal), > PointersArray, 0, > - i); > + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), > + Info.PointersArray, 0, i); > Address PAddr(P, Ctx.getTypeAlignInChars(Ctx.VoidPtrTy)); > CGF.Builder.CreateStore(PVal, PAddr); > > if (hasRuntimeEvaluationCaptureSize) { > llvm::Value *S = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.SizeTy, PointerNumVal), SizesArray, > + llvm::ArrayType::get(CGM.SizeTy, Info.NumberOfPtrs), > + Info.SizesArray, > /*Idx0=*/0, > /*Idx1=*/i); > Address SAddr(S, Ctx.getTypeAlignInChars(Ctx.getSizeType())); > @@ -5682,23 +5800,24 @@ emitOffloadingArrays(CodeGenFunction &CG > static void emitOffloadingArraysArgument( > CodeGenFunction &CGF, llvm::Value *&BasePointersArrayArg, > llvm::Value *&PointersArrayArg, llvm::Value *&SizesArrayArg, > - llvm::Value *&MapTypesArrayArg, llvm::Value *BasePointersArray, > - llvm::Value *PointersArray, llvm::Value *SizesArray, > - llvm::Value *MapTypesArray, unsigned NumElems) { > + llvm::Value *&MapTypesArrayArg, CGOpenMPRuntime::TargetDataInfo &Info) > { > auto &CGM = CGF.CGM; > - if (NumElems) { > + if (Info.NumberOfPtrs) { > BasePointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.VoidPtrTy, NumElems), BasePointersArray, > + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), > + Info.BasePointersArray, > /*Idx0=*/0, /*Idx1=*/0); > PointersArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.VoidPtrTy, NumElems), PointersArray, > + llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs), > + Info.PointersArray, > /*Idx0=*/0, > /*Idx1=*/0); > SizesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.SizeTy, NumElems), SizesArray, > + llvm::ArrayType::get(CGM.SizeTy, Info.NumberOfPtrs), > Info.SizesArray, > /*Idx0=*/0, /*Idx1=*/0); > MapTypesArrayArg = CGF.Builder.CreateConstInBoundsGEP2_32( > - llvm::ArrayType::get(CGM.Int32Ty, NumElems), MapTypesArray, > + llvm::ArrayType::get(CGM.Int32Ty, Info.NumberOfPtrs), > + Info.MapTypesArray, > /*Idx0=*/0, > /*Idx1=*/0); > } else { > @@ -5725,12 +5844,12 @@ void CGOpenMPRuntime::emitTargetCall(Cod > > // Fill up the arrays with all the captured variables. > MappableExprsHandler::MapValuesArrayTy KernelArgs; > - MappableExprsHandler::MapValuesArrayTy BasePointers; > + MappableExprsHandler::MapBaseValuesArrayTy BasePointers; > MappableExprsHandler::MapValuesArrayTy Pointers; > MappableExprsHandler::MapValuesArrayTy Sizes; > MappableExprsHandler::MapFlagsArrayTy MapTypes; > > - MappableExprsHandler::MapValuesArrayTy CurBasePointers; > + MappableExprsHandler::MapBaseValuesArrayTy CurBasePointers; > MappableExprsHandler::MapValuesArrayTy CurPointers; > MappableExprsHandler::MapValuesArrayTy CurSizes; > MappableExprsHandler::MapFlagsArrayTy CurMapTypes; > @@ -5779,7 +5898,7 @@ void CGOpenMPRuntime::emitTargetCall(Cod > > // The kernel args are always the first elements of the base pointers > // associated with a capture. > - KernelArgs.push_back(CurBasePointers.front()); > + KernelArgs.push_back(*CurBasePointers.front()); > // We need to append the results of this capture to what we already > have. > BasePointers.append(CurBasePointers.begin(), CurBasePointers.end()); > Pointers.append(CurPointers.begin(), CurPointers.end()); > @@ -5802,17 +5921,11 @@ void CGOpenMPRuntime::emitTargetCall(Cod > &D](CodeGenFunction &CGF, PrePostActionTy &) { > auto &RT = CGF.CGM.getOpenMPRuntime(); > // Emit the offloading arrays. > - llvm::Value *BasePointersArray; > - llvm::Value *PointersArray; > - llvm::Value *SizesArray; > - llvm::Value *MapTypesArray; > - emitOffloadingArrays(CGF, BasePointersArray, PointersArray, SizesArray, > - MapTypesArray, BasePointers, Pointers, Sizes, > - MapTypes); > - emitOffloadingArraysArgument(CGF, BasePointersArray, PointersArray, > - SizesArray, MapTypesArray, > BasePointersArray, > - PointersArray, SizesArray, MapTypesArray, > - BasePointers.size()); > + TargetDataInfo Info; > + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, > Info); > + emitOffloadingArraysArgument(CGF, Info.BasePointersArray, > + Info.PointersArray, Info.SizesArray, > + Info.MapTypesArray, Info); > > // On top of the arrays that were filled up, the target offloading call > // takes as arguments the device id as well as the host pointer. The > host > @@ -5853,15 +5966,19 @@ void CGOpenMPRuntime::emitTargetCall(Cod > assert(ThreadLimit && "Thread limit expression should be available > along " > "with number of teams."); > llvm::Value *OffloadingArgs[] = { > - DeviceID, OutlinedFnID, PointerNum, > - BasePointersArray, PointersArray, SizesArray, > - MapTypesArray, NumTeams, ThreadLimit}; > + DeviceID, OutlinedFnID, > + PointerNum, Info.BasePointersArray, > + Info.PointersArray, Info.SizesArray, > + Info.MapTypesArray, NumTeams, > + ThreadLimit}; > Return = CGF.EmitRuntimeCall( > RT.createRuntimeFunction(OMPRTL__tgt_target_teams), > OffloadingArgs); > } else { > llvm::Value *OffloadingArgs[] = { > - DeviceID, OutlinedFnID, PointerNum, BasePointersArray, > - PointersArray, SizesArray, MapTypesArray}; > + DeviceID, OutlinedFnID, > + PointerNum, Info.BasePointersArray, > + Info.PointersArray, Info.SizesArray, > + Info.MapTypesArray}; > Return = > CGF.EmitRuntimeCall(RT.createRuntimeFunction(OMPRTL__tgt_target), > OffloadingArgs); > } > @@ -6073,29 +6190,23 @@ void CGOpenMPRuntime::emitNumTeamsClause > PushNumTeamsArgs); > } > > -void CGOpenMPRuntime::emitTargetDataCalls(CodeGenFunction &CGF, > - const OMPExecutableDirective &D, > - const Expr *IfCond, > - const Expr *Device, > - const RegionCodeGenTy &CodeGen) { > - > +void CGOpenMPRuntime::emitTargetDataCalls( > + CodeGenFunction &CGF, const OMPExecutableDirective &D, const Expr > *IfCond, > + const Expr *Device, const RegionCodeGenTy &CodeGen, TargetDataInfo > &Info) { > if (!CGF.HaveInsertPoint()) > return; > > - llvm::Value *BasePointersArray = nullptr; > - llvm::Value *PointersArray = nullptr; > - llvm::Value *SizesArray = nullptr; > - llvm::Value *MapTypesArray = nullptr; > - unsigned NumOfPtrs = 0; > + // Action used to replace the default codegen action and turn > privatization > + // off. > + PrePostActionTy NoPrivAction; > > // Generate the code for the opening of the data environment. Capture all > the > // arguments of the runtime call by reference because they are used in > the > // closing of the region. > - auto &&BeginThenGen = [&D, &CGF, &BasePointersArray, &PointersArray, > - &SizesArray, &MapTypesArray, Device, > - &NumOfPtrs](CodeGenFunction &CGF, PrePostActionTy > &) { > + auto &&BeginThenGen = [&D, &CGF, Device, &Info, &CodeGen, &NoPrivAction]( > + CodeGenFunction &CGF, PrePostActionTy &) { > // Fill up the arrays with all the mapped variables. > - MappableExprsHandler::MapValuesArrayTy BasePointers; > + MappableExprsHandler::MapBaseValuesArrayTy BasePointers; > MappableExprsHandler::MapValuesArrayTy Pointers; > MappableExprsHandler::MapValuesArrayTy Sizes; > MappableExprsHandler::MapFlagsArrayTy MapTypes; > @@ -6103,21 +6214,16 @@ void CGOpenMPRuntime::emitTargetDataCall > // Get map clause information. > MappableExprsHandler MCHandler(D, CGF); > MCHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); > - NumOfPtrs = BasePointers.size(); > > // Fill up the arrays and create the arguments. > - emitOffloadingArrays(CGF, BasePointersArray, PointersArray, SizesArray, > - MapTypesArray, BasePointers, Pointers, Sizes, > - MapTypes); > + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, > Info); > > llvm::Value *BasePointersArrayArg = nullptr; > llvm::Value *PointersArrayArg = nullptr; > llvm::Value *SizesArrayArg = nullptr; > llvm::Value *MapTypesArrayArg = nullptr; > emitOffloadingArraysArgument(CGF, BasePointersArrayArg, > PointersArrayArg, > - SizesArrayArg, MapTypesArrayArg, > - BasePointersArray, PointersArray, > SizesArray, > - MapTypesArray, NumOfPtrs); > + SizesArrayArg, MapTypesArrayArg, Info); > > // Emit device ID if any. > llvm::Value *DeviceID = nullptr; > @@ -6128,7 +6234,7 @@ void CGOpenMPRuntime::emitTargetDataCall > DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF); > > // Emit the number of elements in the offloading arrays. > - auto *PointerNum = CGF.Builder.getInt32(NumOfPtrs); > + auto *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs); > > llvm::Value *OffloadingArgs[] = { > DeviceID, PointerNum, BasePointersArrayArg, > @@ -6136,23 +6242,24 @@ void CGOpenMPRuntime::emitTargetDataCall > auto &RT = CGF.CGM.getOpenMPRuntime(); > > CGF.EmitRuntimeCall(RT.createRuntimeFunction(OMPRTL__tgt_target_data_begin), > OffloadingArgs); > + > + // If device pointer privatization is required, emit the body of the > region > + // here. It will have to be duplicated: with and without privatization. > + if (!Info.CaptureDeviceAddrMap.empty()) > + CodeGen(CGF); > }; > > // Generate code for the closing of the data region. > - auto &&EndThenGen = [&CGF, &BasePointersArray, &PointersArray, > &SizesArray, > - &MapTypesArray, Device, > - &NumOfPtrs](CodeGenFunction &CGF, PrePostActionTy &) > { > - assert(BasePointersArray && PointersArray && SizesArray && > MapTypesArray && > - NumOfPtrs && "Invalid data environment closing arguments."); > + auto &&EndThenGen = [&CGF, Device, &Info](CodeGenFunction &CGF, > + PrePostActionTy &) { > + assert(Info.isValid() && "Invalid data environment closing > arguments."); > > llvm::Value *BasePointersArrayArg = nullptr; > llvm::Value *PointersArrayArg = nullptr; > llvm::Value *SizesArrayArg = nullptr; > llvm::Value *MapTypesArrayArg = nullptr; > emitOffloadingArraysArgument(CGF, BasePointersArrayArg, > PointersArrayArg, > - SizesArrayArg, MapTypesArrayArg, > - BasePointersArray, PointersArray, > SizesArray, > - MapTypesArray, NumOfPtrs); > + SizesArrayArg, MapTypesArrayArg, Info); > > // Emit device ID if any. > llvm::Value *DeviceID = nullptr; > @@ -6163,7 +6270,7 @@ void CGOpenMPRuntime::emitTargetDataCall > DeviceID = CGF.Builder.getInt32(OMP_DEVICEID_UNDEF); > > // Emit the number of elements in the offloading arrays. > - auto *PointerNum = CGF.Builder.getInt32(NumOfPtrs); > + auto *PointerNum = CGF.Builder.getInt32(Info.NumberOfPtrs); > > llvm::Value *OffloadingArgs[] = { > DeviceID, PointerNum, BasePointersArrayArg, > @@ -6173,24 +6280,40 @@ void CGOpenMPRuntime::emitTargetDataCall > OffloadingArgs); > }; > > - // In the event we get an if clause, we don't have to take any action on > the > - // else side. > - auto &&ElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {}; > + // If we need device pointer privatization, we need to emit the body of > the > + // region with no privatization in the 'else' branch of the conditional. > + // Otherwise, we don't have to do anything. > + auto &&BeginElseGen = [&Info, &CodeGen, &NoPrivAction](CodeGenFunction > &CGF, > + PrePostActionTy &) > { > + if (!Info.CaptureDeviceAddrMap.empty()) { > + CodeGen.setAction(NoPrivAction); > + CodeGen(CGF); > + } > + }; > + > + // We don't have to do anything to close the region if the if clause > evaluates > + // to false. > + auto &&EndElseGen = [](CodeGenFunction &CGF, PrePostActionTy &) {}; > > if (IfCond) { > - emitOMPIfClause(CGF, IfCond, BeginThenGen, ElseGen); > + emitOMPIfClause(CGF, IfCond, BeginThenGen, BeginElseGen); > } else { > - RegionCodeGenTy BeginThenRCG(BeginThenGen); > - BeginThenRCG(CGF); > + RegionCodeGenTy RCG(BeginThenGen); > + RCG(CGF); > } > > - CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_target_data, > CodeGen); > + // If we don't require privatization of device pointers, we emit the body > in > + // between the runtime calls. This avoids duplicating the body code. > + if (Info.CaptureDeviceAddrMap.empty()) { > + CodeGen.setAction(NoPrivAction); > + CodeGen(CGF); > + } > > if (IfCond) { > - emitOMPIfClause(CGF, IfCond, EndThenGen, ElseGen); > + emitOMPIfClause(CGF, IfCond, EndThenGen, EndElseGen); > } else { > - RegionCodeGenTy EndThenRCG(EndThenGen); > - EndThenRCG(CGF); > + RegionCodeGenTy RCG(EndThenGen); > + RCG(CGF); > } > } > > @@ -6208,7 +6331,7 @@ void CGOpenMPRuntime::emitTargetDataStan > // Generate the code for the opening of the data environment. > auto &&ThenGen = [&D, &CGF, Device](CodeGenFunction &CGF, PrePostActionTy > &) { > // Fill up the arrays with all the mapped variables. > - MappableExprsHandler::MapValuesArrayTy BasePointers; > + MappableExprsHandler::MapBaseValuesArrayTy BasePointers; > MappableExprsHandler::MapValuesArrayTy Pointers; > MappableExprsHandler::MapValuesArrayTy Sizes; > MappableExprsHandler::MapFlagsArrayTy MapTypes; > @@ -6217,19 +6340,12 @@ void CGOpenMPRuntime::emitTargetDataStan > MappableExprsHandler MEHandler(D, CGF); > MEHandler.generateAllInfo(BasePointers, Pointers, Sizes, MapTypes); > > - llvm::Value *BasePointersArrayArg = nullptr; > - llvm::Value *PointersArrayArg = nullptr; > - llvm::Value *SizesArrayArg = nullptr; > - llvm::Value *MapTypesArrayArg = nullptr; > - > // Fill up the arrays and create the arguments. > - emitOffloadingArrays(CGF, BasePointersArrayArg, PointersArrayArg, > - SizesArrayArg, MapTypesArrayArg, BasePointers, > - Pointers, Sizes, MapTypes); > - emitOffloadingArraysArgument( > - CGF, BasePointersArrayArg, PointersArrayArg, SizesArrayArg, > - MapTypesArrayArg, BasePointersArrayArg, PointersArrayArg, > SizesArrayArg, > - MapTypesArrayArg, BasePointers.size()); > + TargetDataInfo Info; > + emitOffloadingArrays(CGF, BasePointers, Pointers, Sizes, MapTypes, > Info); > + emitOffloadingArraysArgument(CGF, Info.BasePointersArray, > + Info.PointersArray, Info.SizesArray, > + Info.MapTypesArray, Info); > > // Emit device ID if any. > llvm::Value *DeviceID = nullptr; > @@ -6243,8 +6359,8 @@ void CGOpenMPRuntime::emitTargetDataStan > auto *PointerNum = CGF.Builder.getInt32(BasePointers.size()); > > llvm::Value *OffloadingArgs[] = { > - DeviceID, PointerNum, BasePointersArrayArg, > - PointersArrayArg, SizesArrayArg, MapTypesArrayArg}; > + DeviceID, PointerNum, Info.BasePointersArray, > + Info.PointersArray, Info.SizesArray, Info.MapTypesArray}; > > auto &RT = CGF.CGM.getOpenMPRuntime(); > // Select the right runtime function call for each expected standalone > > Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original) > +++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Thu Jul 28 09:23:26 2016 > @@ -997,17 +997,59 @@ public: > virtual void emitNumTeamsClause(CodeGenFunction &CGF, const Expr > *NumTeams, > const Expr *ThreadLimit, SourceLocation > Loc); > > + /// Struct that keeps all the relevant information that should be kept > + /// throughout a 'target data' region. > + class TargetDataInfo { > + /// Set to true if device pointer information have to be obtained. > + bool RequiresDevicePointerInfo = false; > + > + public: > + /// The array of base pointer passed to the runtime library. > + llvm::Value *BasePointersArray = nullptr; > + /// The array of section pointers passed to the runtime library. > + llvm::Value *PointersArray = nullptr; > + /// The array of sizes passed to the runtime library. > + llvm::Value *SizesArray = nullptr; > + /// The array of map types passed to the runtime library. > + llvm::Value *MapTypesArray = nullptr; > + /// The total number of pointers passed to the runtime library. > + unsigned NumberOfPtrs = 0u; > + /// Map between the a declaration of a capture and the corresponding > base > + /// pointer address where the runtime returns the device pointers. > + llvm::DenseMap<const ValueDecl *, Address> CaptureDeviceAddrMap; > + > + explicit TargetDataInfo() {} > + explicit TargetDataInfo(bool RequiresDevicePointerInfo) > + : RequiresDevicePointerInfo(RequiresDevicePointerInfo) {} > + /// Clear information about the data arrays. > + void clearArrayInfo() { > + BasePointersArray = nullptr; > + PointersArray = nullptr; > + SizesArray = nullptr; > + MapTypesArray = nullptr; > + NumberOfPtrs = 0u; > + } > + /// Return true if the current target data information has valid > arrays. > + bool isValid() { > + return BasePointersArray && PointersArray && SizesArray && > + MapTypesArray && NumberOfPtrs; > + } > + bool requiresDevicePointerInfo() { return RequiresDevicePointerInfo; } > + }; > + > /// \brief Emit the target data mapping code associated with \a D. > /// \param D Directive to emit. > - /// \param IfCond Expression evaluated in if clause associated with the > target > - /// directive, or null if no if clause is used. > + /// \param IfCond Expression evaluated in if clause associated with the > + /// target directive, or null if no device clause is used. > /// \param Device Expression evaluated in device clause associated with > the > /// target directive, or null if no device clause is used. > - /// \param CodeGen Function that emits the enclosed region. > + /// \param Info A record used to store information that needs to be > preserved > + /// until the region is closed. > virtual void emitTargetDataCalls(CodeGenFunction &CGF, > const OMPExecutableDirective &D, > const Expr *IfCond, const Expr *Device, > - const RegionCodeGenTy &CodeGen); > + const RegionCodeGenTy &CodeGen, > + TargetDataInfo &Info); > > /// \brief Emit the data mapping/movement code associated with the > directive > /// \a D that should be of the form 'target [{enter|exit} data | > update]'. > > Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original) > +++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Thu Jul 28 09:23:26 2016 > @@ -3400,22 +3400,137 @@ CodeGenFunction::getOMPCancelDestination > return BreakContinueStack.back().BreakBlock; > } > > +void CodeGenFunction::EmitOMPUseDevicePtrClause( > + const OMPClause &NC, OMPPrivateScope &PrivateScope, > + const llvm::DenseMap<const ValueDecl *, Address> &CaptureDeviceAddrMap) > { > + const auto &C = cast<OMPUseDevicePtrClause>(NC); > + auto OrigVarIt = C.varlist_begin(); > + auto InitIt = C.inits().begin(); > + for (auto PvtVarIt : C.private_copies()) { > + auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(*OrigVarIt)->getDecl()); > + auto *InitVD = cast<VarDecl>(cast<DeclRefExpr>(*InitIt)->getDecl()); > + auto *PvtVD = cast<VarDecl>(cast<DeclRefExpr>(PvtVarIt)->getDecl()); > + > + // In order to identify the right initializer we need to match the > + // declaration used by the mapping logic. In some cases we may get > + // OMPCapturedExprDecl that refers to the original declaration. > + const ValueDecl *MatchingVD = OrigVD; > + if (auto *OED = dyn_cast<OMPCapturedExprDecl>(MatchingVD)) { > + // OMPCapturedExprDecl are used to privative fields of the current > + // structure. > + auto *ME = cast<MemberExpr>(OED->getInit()); > + assert(isa<CXXThisExpr>(ME->getBase()) && > + "Base should be the current struct!"); > + MatchingVD = ME->getMemberDecl(); > + } > + > + // If we don't have information about the current list item, move on to > + // the next one. > + auto InitAddrIt = CaptureDeviceAddrMap.find(MatchingVD); > + if (InitAddrIt == CaptureDeviceAddrMap.end()) > + continue; > + > + bool IsRegistered = PrivateScope.addPrivate(OrigVD, [&]() -> Address { > + // Initialize the temporary initialization variable with the address > we > + // get from the runtime library. We have to cast the source address > + // because it is always a void *. References are materialized in the > + // privatization scope, so the initialization here disregards the > fact > + // the original variable is a reference. > + QualType AddrQTy = > + > getContext().getPointerType(OrigVD->getType().getNonReferenceType()); > + llvm::Type *AddrTy = ConvertTypeForMem(AddrQTy); > + Address InitAddr = Builder.CreateBitCast(InitAddrIt->second, AddrTy); > + setAddrOfLocalVar(InitVD, InitAddr); > + > + // Emit private declaration, it will be initialized by the value we > + // declaration we just added to the local declarations map. > + EmitDecl(*PvtVD); > + > + // The initialization variables reached its purpose in the emission > + // ofthe previous declaration, so we don't need it anymore. > + LocalDeclMap.erase(InitVD); > + > + // Return the address of the private variable. > + return GetAddrOfLocalVar(PvtVD); > + }); > + assert(IsRegistered && "firstprivate var already registered as > private"); > + // Silence the warning about unused variable. > + (void)IsRegistered; > + > + ++OrigVarIt; > + ++InitIt; > + } > +} > + > // Generate the instructions for '#pragma omp target data' directive. > void CodeGenFunction::EmitOMPTargetDataDirective( > const OMPTargetDataDirective &S) { > - // The target data enclosed region is implemented just by emitting the > - // statement. > - auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) { > - > CGF.EmitStmt(cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt()); > + CGOpenMPRuntime::TargetDataInfo Info(/*RequiresDevicePointerInfo=*/true); > + > + // Create a pre/post action to signal the privatization of the device > pointer. > + // This action can be replaced by the OpenMP runtime code generation to > + // deactivate privatization. > + bool PrivatizeDevicePointers = false; > + class DevicePointerPrivActionTy : public PrePostActionTy { > + bool &PrivatizeDevicePointers; > + > + public: > + explicit DevicePointerPrivActionTy(bool &PrivatizeDevicePointers) > + : PrePostActionTy(), > PrivatizeDevicePointers(PrivatizeDevicePointers) {} > + void Enter(CodeGenFunction &CGF) override { > + PrivatizeDevicePointers = true; > + } > }; > + DevicePointerPrivActionTy PrivAction(PrivatizeDevicePointers); > + > + auto &&CodeGen = [&S, &Info, &PrivatizeDevicePointers]( > + CodeGenFunction &CGF, PrePostActionTy &Action) { > + auto &&InnermostCodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &) > { > + CGF.EmitStmt( > + cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt()); > + }; > + > + // Codegen that selects wheather to generate the privatization code or > not. > + auto &&PrivCodeGen = [&S, &Info, &PrivatizeDevicePointers, > + &InnermostCodeGen](CodeGenFunction &CGF, > + PrePostActionTy &Action) { > + RegionCodeGenTy RCG(InnermostCodeGen); > + PrivatizeDevicePointers = false; > + > + // Call the pre-action to change the status of > PrivatizeDevicePointers if > + // needed. > + Action.Enter(CGF); > + > + if (PrivatizeDevicePointers) { > + OMPPrivateScope PrivateScope(CGF); > + // Emit all instances of the use_device_ptr clause. > + for (const auto *C : S.getClausesOfKind<OMPUseDevicePtrClause>()) > + CGF.EmitOMPUseDevicePtrClause(*C, PrivateScope, > + Info.CaptureDeviceAddrMap); > + (void)PrivateScope.Privatize(); > + RCG(CGF); > + } else > + RCG(CGF); > + }; > + > + // Forward the provided action to the privatization codegen. > + RegionCodeGenTy PrivRCG(PrivCodeGen); > + PrivRCG.setAction(Action); > + > + // Notwithstanding the body of the region is emitted as inlined > directive, > + // we don't use an inline scope as changes in the references inside the > + // region are expected to be visible outside, so we do not privative > them. > + OMPLexicalScope Scope(CGF, S); > + CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_target_data, > + PrivRCG); > + }; > + > + RegionCodeGenTy RCG(CodeGen); > > // If we don't have target devices, don't bother emitting the data > mapping > // code. > if (CGM.getLangOpts().OMPTargetTriples.empty()) { > - OMPLexicalScope Scope(*this, S, /*AsInlined=*/true); > - > - CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_target_data, > - CodeGen); > + RCG(*this); > return; > } > > @@ -3429,7 +3544,12 @@ void CodeGenFunction::EmitOMPTargetDataD > if (auto *C = S.getSingleClause<OMPDeviceClause>()) > Device = C->getDevice(); > > - CGM.getOpenMPRuntime().emitTargetDataCalls(*this, S, IfCond, Device, > CodeGen); > + // Set the action to signal privatization of device pointers. > + RCG.setAction(PrivAction); > + > + // Emit region code. > + CGM.getOpenMPRuntime().emitTargetDataCalls(*this, S, IfCond, Device, RCG, > + Info); > } > > void CodeGenFunction::EmitOMPTargetEnterDataDirective( > > Modified: cfe/trunk/lib/CodeGen/CodeGenFunction.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenFunction.h?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/CodeGen/CodeGenFunction.h (original) > +++ cfe/trunk/lib/CodeGen/CodeGenFunction.h Thu Jul 28 09:23:26 2016 > @@ -2392,6 +2392,9 @@ public: > OMPPrivateScope &PrivateScope); > void EmitOMPPrivateClause(const OMPExecutableDirective &D, > OMPPrivateScope &PrivateScope); > + void EmitOMPUseDevicePtrClause( > + const OMPClause &C, OMPPrivateScope &PrivateScope, > + const llvm::DenseMap<const ValueDecl *, Address> > &CaptureDeviceAddrMap); > /// \brief Emit code for copyin clause in \a D directive. The next code > is > /// generated at the start of outlined functions for directives: > /// \code > > Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original) > +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Thu Jul 28 09:23:26 2016 > @@ -11800,7 +11800,10 @@ OMPClause *Sema::ActOnOpenMPUseDevicePtr > SourceLocation StartLoc, > SourceLocation LParenLoc, > SourceLocation EndLoc) { > - SmallVector<Expr *, 8> Vars; > + MappableVarListInfo MVLI(VarList); > + SmallVector<Expr *, 8> PrivateCopies; > + SmallVector<Expr *, 8> Inits; > + > for (auto &RefExpr : VarList) { > assert(RefExpr && "NULL expr in OpenMP use_device_ptr clause."); > SourceLocation ELoc; > @@ -11809,27 +11812,73 @@ OMPClause *Sema::ActOnOpenMPUseDevicePtr > auto Res = getPrivateItem(*this, SimpleRefExpr, ELoc, ERange); > if (Res.second) { > // It will be analyzed later. > - Vars.push_back(RefExpr); > + MVLI.ProcessedVarList.push_back(RefExpr); > + PrivateCopies.push_back(nullptr); > + Inits.push_back(nullptr); > } > ValueDecl *D = Res.first; > if (!D) > continue; > > QualType Type = D->getType(); > - // item should be a pointer or reference to pointer > - if (!Type.getNonReferenceType()->isPointerType()) { > + Type = Type.getNonReferenceType().getUnqualifiedType(); > + > + auto *VD = dyn_cast<VarDecl>(D); > + > + // Item should be a pointer or reference to pointer. > + if (!Type->isPointerType()) { > Diag(ELoc, diag::err_omp_usedeviceptr_not_a_pointer) > << 0 << RefExpr->getSourceRange(); > continue; > } > - Vars.push_back(RefExpr->IgnoreParens()); > + > + // Build the private variable and the expression that refers to it. > + auto VDPrivate = buildVarDecl(*this, ELoc, Type, D->getName(), > + D->hasAttrs() ? &D->getAttrs() : > nullptr); > + if (VDPrivate->isInvalidDecl()) > + continue; > + > + CurContext->addDecl(VDPrivate); > + auto VDPrivateRefExpr = buildDeclRefExpr( > + *this, VDPrivate, RefExpr->getType().getUnqualifiedType(), ELoc); > + > + // Add temporary variable to initialize the private copy of the > pointer. > + auto *VDInit = > + buildVarDecl(*this, RefExpr->getExprLoc(), Type, ".devptr.temp"); > + auto *VDInitRefExpr = buildDeclRefExpr(*this, VDInit, > RefExpr->getType(), > + RefExpr->getExprLoc()); > + AddInitializerToDecl(VDPrivate, > + DefaultLvalueConversion(VDInitRefExpr).get(), > + /*DirectInit=*/false, > /*TypeMayContainAuto=*/false); > + > + // If required, build a capture to implement the privatization > initialized > + // with the current list item value. > + DeclRefExpr *Ref = nullptr; > + if (!VD) > + Ref = buildCapture(*this, D, SimpleRefExpr, /*WithInit=*/true); > + MVLI.ProcessedVarList.push_back(VD ? RefExpr->IgnoreParens() : Ref); > + PrivateCopies.push_back(VDPrivateRefExpr); > + Inits.push_back(VDInitRefExpr); > + > + // We need to add a data sharing attribute for this variable to make > sure it > + // is correctly captured. A variable that shows up in a use_device_ptr > has > + // similar properties of a first private variable. > + DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref); > + > + // Create a mappable component for the list item. List items in this > clause > + // only need a component. > + MVLI.VarBaseDeclarations.push_back(D); > + MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); > + MVLI.VarComponents.back().push_back( > + OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D)); > } > > - if (Vars.empty()) > + if (MVLI.ProcessedVarList.empty()) > return nullptr; > > - return OMPUseDevicePtrClause::Create(Context, StartLoc, LParenLoc, > EndLoc, > - Vars); > + return OMPUseDevicePtrClause::Create( > + Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList, > + PrivateCopies, Inits, MVLI.VarBaseDeclarations, MVLI.VarComponents); > } > > OMPClause *Sema::ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList, > > Modified: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=276977&r1=276976&r2=276977&view=diff > ============================================================================== > --- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp (original) > +++ cfe/tr > > -- > Mike > Sent from phone > > > > > _______________________________________________ > cfe-commits mailing list > cfe-commits@lists.llvm.org > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits > _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits