I agree this looks like a compiler bug, but I'm not sure what's causing it exactly. Putting "this->" in front of those accesses to CurDir and CGF seems to fix it, but it's also very ugly.
I tried a couple of other things, like moving the method out-of-line, but no luck yet. On Fri, Jul 29, 2016 at 1:45 PM, Samuel F Antao <sfan...@us.ibm.com> wrote: > I don't have easy access to Windows machine to test this, but it seems that > whichever compiler is being used is broken. This is about a member function > accessing a variable of the same class. I don't think I am doing any > non-portable code pattern here. > > Alexey, do you have a easy way to test this? Can you give me a little help > here? > > Thanks! > Samuel > > > ----- Original message ----- > From: Hans Wennborg <h...@chromium.org> > Sent by: hwennb...@google.com > To: Samuel F Antao/Watson/IBM@IBMUS > Cc: Mike Aizatsky <aizat...@google.com>, cfe-commits > <cfe-commits@lists.llvm.org> > Subject: Re: r276977 - [OpenMP] Codegen for use_device_ptr clause. > Date: Fri, Jul 29, 2016 3:53 PM > > 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
