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